	.global .u64 %_assertInfo;

.func assertCtor(.reg .u64 assertInfo)
{
	/*
	.reg .v4 .u16 tid;
	.reg .pred q;

	mov.v4.b16 tid, %tid;
	setp.eq.u16 q, tid.x, 0;
@!q	setp.eq.u16 q, tid.y, 0;
@!q	setp.eq.u16 q, tid.z, 10U;
@!q	ret;
	*/
	st.volatile.global.u64 [%_assertInfo], assertInfo;
	st.global.u32 [assertInfo+0], 0U;
	st.global.u32 [assertInfo+4], 0U;
	ret;
}
// Utility Methods
.func assertFail(.reg .b32 assertId) { .reg .u64 %r<1>; ld.global.u64 %r0, [%_assertInfo]; st.global.u32 [%r0+0], assertId; exit; }
// Equality Asserts
.func assertAreEqual_b64(.reg .b32 assertId, .reg .b64 expected, .reg .b64 actual) { .reg .pred %p; setp.eq.b64 %p, expected, actual; @%p ret; call assertFail, (assertId); }
.func assertAreNotEqual_b64(.reg .b32 assertId, .reg .b64 expected, .reg .b64 actual) { .reg .pred %p; setp.ne.b64 %p, expected, actual; @%p ret; call assertFail, (assertId); }
// Comparison Asserts
.func assertGreater_s64(.reg .b32 assertId, .reg .s64 arg1, .reg .s64 arg2) { .reg .pred %p; setp.gt.s64 %p, arg1, arg2; @%p ret; call assertFail, (assertId); }
.func assertGreater_u64(.reg .b32 assertId, .reg .u64 arg1, .reg .u64 arg2) { .reg .pred %p; setp.hi.u64 %p, arg1, arg2; @%p ret; call assertFail, (assertId); }
.func assertGreaterOrEqual_s64(.reg .b32 assertId, .reg .s64 arg1, .reg .s64 arg2) { .reg .pred %p; setp.ge.s64 %p, arg1, arg2; @%p ret; call assertFail, (assertId); }
.func assertGreaterOrEqual_u64(.reg .b32 assertId, .reg .u64 arg1, .reg .u64 arg2) { .reg .pred %p; setp.hs.u64 %p, arg1, arg2; @%p ret; call assertFail, (assertId); }
.func assertLess_s64(.reg .b32 assertId, .reg .s64 arg1, .reg .s64 arg2) { .reg .pred %p; setp.lt.u64 %p, arg1, arg2; @%p ret; call assertFail, (assertId); }
.func assertLess_u64(.reg .b32 assertId, .reg .u64 arg1, .reg .u64 arg2) { .reg .pred %p; setp.lo.u64 %p, arg1, arg2; @%p ret; call assertFail, (assertId); }
.func assertLessOrEqual_s64(.reg .b32 assertId, .reg .s64 arg1, .reg .s64 arg2) { .reg .pred %p; setp.le.u64 %p, arg1, arg2; @%p ret; call assertFail, (assertId); }
.func assertLessOrEqual_u64(.reg .b32 assertId, .reg .u64 arg1, .reg .u64 arg2) { .reg .pred %p; setp.ls.u64 %p, arg1, arg2; @%p ret; call assertFail, (assertId); }
// Condition Tests
.func assertIsTrue_pred(.reg .b32 assertId, .reg .pred condition) { @condition ret; call assertFail, (assertId); }
.func assertIsFalse_pred(.reg .b32 assertId, .reg .pred condition) { @!condition ret; call assertFail, (assertId); }
.func assertIsNull_b64(.reg .b32 assertId, .reg .b64 object) { .reg .pred %p; setp.eq.u64 %p, object, 0U; @%p ret; call assertFail, (assertId); }
.func assertIsNotNull_b64(.reg .b32 assertId, .reg .b64 object) { .reg .pred %p; setp.eq.u64 %p, object, 0U; @!%p ret; call assertFail, (assertId); }
